#   Copyright 2010 Matthew Hawn

#   Licensed under the Apache License, Version 2.0 (the "License");
#   you may not use this file except in compliance with the License.
#   You may obtain a copy of the License at

#       http://www.apache.org/licenses/LICENSE-2.0

#   Unless required by applicable law or agreed to in writing, software
#   distributed under the License is distributed on an "AS IS" BASIS,
#   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
#   See the License for the specific language governing permissions and
#   limitations under the License.

#all the external stuff needed
from stdlib cimport *
from python_sequence cimport PySequence_Check
from python_string cimport PyString_FromStringAndSize, PyString_AsString, PyString_AsStringAndSize
cimport cython
from python_dict cimport *
from python_long cimport PyLong_FromVoidPtr
from python_buffer cimport (
    PyObject_GetBuffer,
    PyObject_CheckBuffer, 
    PyBuffer_IsContiguous,
    PyBUF_SIMPLE, 
    PyBUF_WRITABLE, 
    PyBUF_ND,
    PyBUF_STRIDES,
    PyBUF_FORMAT,
    PyBUF_C_CONTIGUOUS,
    PyBUF_F_CONTIGUOUS,
    PyBUF_ANY_CONTIGUOUS )
from python_mem cimport *


cdef extern from "Python.h":
        void PyBuffer_Release(Py_buffer * view) #bug in cython 0.12
        int PyBuffer_FillInfo(Py_buffer *view, obj, void *buf,
                          Py_ssize_t len, int readonly,
                          int flags) except -1
        int PyObject_AsWriteBuffer(obj, void **buffer,  Py_ssize_t *buffer_len) except -1
        int PyObject_AsReadBuffer(obj, void **buffer,  Py_ssize_t *buffer_len) except -1
        void PySys_WriteStderr(char *format, ...)
        int Py_EQ
        int Py_NE
cdef extern from "string.h" nogil:
    void *memset(void *, int, size_t)

#predefines
cdef class ErrorList
cdef class ConstantList
cdef class BitfieldList(ConstantList)
cdef class IntConstantList(ConstantList)
cdef class UIntConstantList(IntConstantList)
cdef class Constant
cdef class InfoGetter
cdef class Referenced(InfoGetter)
cdef class Platform(InfoGetter)
cdef class Device(InfoGetter)
cdef class Context(Referenced)
cdef class CommandQueue(Referenced) #TODO: thread-safety
cdef class MemObject(Referenced)
cdef class Image(MemObject)
cdef class Buffer(MemObject)
cdef class Event(Referenced)
cdef class Sampler(Referenced)
cdef class Image2D(Image)
cdef class Image3D(Image)
cdef class Program(Referenced)
cdef class MemMap
cdef class Kernel(Referenced)
cdef class LocalMemory

#Include definitions for opencl functions and constants
# Note: Generated by generate.py from CL/cl.h
ctypedef void * void_ptr
include "cl.pxi" 

#random function typedefs
ctypedef cl_int (*InfoFunc)(void_ptr, cl_uint, size_t, void *, size_t *)
ctypedef cl_int (*ExtraInfoFunc) (void_ptr, void_ptr, cl_uint, size_t, void *, size_t *)
ctypedef object (*ConstFactory)( parent, value)
ctypedef object (*CreateFunc)(void *,object)
ctypedef cl_int (*ReferenceFunc)(void_ptr)

ctypedef union u_value:
    cl_uint v32
    cl_ulong v64

#global def
DEF MAX_DIM=3

#-------------------------------------
# Helper functions
#-----------------------------------------

#returns allocated waitlist array or NULL. Caller must free with PyMem_Free
cdef cl_uint create_events(object events, cl_event **ret_wait) except? 0:
    cdef tuple events_tuple 
    cdef Event e
    cdef cl_event * c_wait = NULL
    cdef cl_uint size = 0
    
    if events is not None:
        size = len(events_tuple)
        c_wait = <cl_event *> PyMem_Malloc(sizeof(cl_event) * size)
        try:
            for i,e in enumerate(events):
                c_wait[i] = e.id
        except:
            PyMem_Free(c_wait)
            raise
            
    ret_wait[0] = c_wait
    return size


#fill out origin and region arrays from sequences
cdef get_image_params(int dim, cl_mem id, origin, region, Py_buffer *default_view, 
        size_t *c_origin, size_t *c_region):
    
    
    memset(c_origin,0, sizeof(size_t) * 3) 
    if origin is not None:
        if not PySequence_Check(origin) or len(origin) != dim:
            raise TypeError("Origin must be a sequence of %d length" %(dim,))
        
        for i in range(dim):
            c_origin[i] = origin[i]

    c_region[2] =1    
    if default_view != NULL:  
        for i in range(dim):
            c_region[i] = default_view.shape[i]   
    elif region is None:
        #pull defaults from image
        check_error("Getting Image Info", 
            clGetImageInfo(id, CL_IMAGE_WIDTH, sizeof(size_t), c_region, NULL))
        check_error("Getting Image Info", 
            clGetImageInfo(id, CL_IMAGE_HEIGHT, sizeof(size_t), c_region+1, NULL))    
        if dim == 3:
            check_error("Getting Image Info", 
                clGetImageInfo(id, CL_IMAGE_DEPTH, sizeof(size_t), c_region+2, NULL))
    else:
        if not PySequence_Check(region) or len(region) != dim:
            raise TypeError("Region must be a sequence of %d length" %(dim,))
        for i in range(dim):
            c_region[i] = region[i]
                
#fast dict reversal (cython optimizes this pretty well)
cdef dict reverse_dict(dict d):
    cdef dict out = dict()
    for k,v in d.iteritems():
        out[v] = k
    return out
    
#generalized function for retreiving buffer using *new* buffer interface (prefered),
#   or older buffer interface.  
# Must call release_buffer when done
#TODO: make this python 3 compatible
cdef Py_ssize_t get_buffer(data, bint writable, void**ptr, Py_buffer *view, int flags = PyBUF_SIMPLE) except? 0:
    cdef Py_ssize_t l
    view.buf = NULL
    if PyObject_CheckBuffer(data):
        if writable:
            flags |= PyBUF_WRITABLE
        
        PyObject_GetBuffer(data, view, flags)
        ptr[0] = view.buf
        return view.len
    else:
        #older interface
        if writable:
            PyObject_AsWriteBuffer(data, ptr, &l)
        else:
            PyObject_AsReadBuffer(data,ptr, &l)
    
        return l

#release above buffer
cdef release_buffer(Py_buffer *view):
    if view.buf != NULL:
        PyBuffer_Release(view)


cdef cl_platform_id get_first_platform() except NULL:
    cdef cl_platform_id platform
    
    check_error("Get Platforms", clGetPlatformIDs(1, &platform,NULL))
    return platform
    
cdef dict c_image_format_channel_dict
cdef dict c_image_format_type_dict
cdef dict c_image_format_size_dict
c_image_format_channel_dict = {
    CL_R            : 1,
    CL_A            : 1,
    CL_INTENSITY    : 1,
    CL_LUMINANCE    : 1,
    CL_RG           : 2,
    CL_RA           : 2,
    CL_RGB          : 1,   #packed format
    CL_RGBA         : 4,
    CL_ARGB         : 4,
    CL_BGRA         : 4,
}
c_image_format_type_dict = {
    CL_SNORM_INT8: 'b',
    CL_SNORM_INT16: 'h',
    CL_UNORM_INT8: 'B',
    CL_UNORM_INT16: 'H',
    CL_UNORM_SHORT_565: 'H',
    CL_UNORM_SHORT_555: 'H',
    CL_UNORM_INT_101010: 'I',
    CL_SIGNED_INT8: 'b',
    CL_SIGNED_INT16: 'h',
    CL_SIGNED_INT32: 'i',
    CL_UNSIGNED_INT8: 'B',
    CL_UNSIGNED_INT16: 'H',
    CL_UNSIGNED_INT32: 'I',
    CL_HALF_FLOAT: 'H',
    CL_FLOAT: 'f',
}
c_image_format_size_dict = {
    CL_SNORM_INT8: 1,
    CL_SNORM_INT16: 2,
    CL_UNORM_INT8: 1,
    CL_UNORM_INT16: 2,
    CL_UNORM_SHORT_565: 2,
    CL_UNORM_SHORT_555: 2,
    CL_UNORM_INT_101010: 4,
    CL_SIGNED_INT8: 1,
    CL_SIGNED_INT16: 2,
    CL_SIGNED_INT32: 4,
    CL_UNSIGNED_INT8: 1,
    CL_UNSIGNED_INT16: 2,
    CL_UNSIGNED_INT32: 4,
    CL_HALF_FLOAT: 2,
    CL_FLOAT: 4,
}    
    
cdef size_t calc_image_format(cl_image_format *format, char *format_str  ):
    cdef int channels = c_image_format_channel_dict[format.image_channel_order]
    cdef int size = c_image_format_size_dict[format.image_channel_data_type]
    cdef char data_type = c_image_format_type_dict[format.image_channel_data_type]
    
    format_str[0] = '='
    format_str += 1
    for i in range(channels):
        format_str[0] = data_type
        format_str += 1
    format_str[0] = 0
    return channels * size

#----------------------------------------------
# Error Handlinng
#-------------------------------------------------

cdef inline check_error(context, cl_int err):
    global c_cl_error
    if err != CL_SUCCESS:
        raise c_cl_error[err](context)
        
cdef inline raise_error(context, cl_int err):
    global c_cl_error
    raise c_cl_error[err](context)

cdef class ErrorList:
    cdef dict mappings
    cdef dict reverse
    cdef dict cache
    cdef str name
    
    cpdef as_string(self, err):
        cdef PyObject *o
        o = PyDict_GetItem(self.reverse, err)
        if o == NULL:
            return ("Undefined Error: %d" % (err))
        return  <object>o  # returning new reference
    
    def __getattr__(self, attr):
        cdef PyObject *o
        o = PyDict_GetItem(self.cache, attr)
        if o != NULL:
            return <object> o #return new reference
    
        o = PyDict_GetItem(self.mappings, attr)
        if o != NULL:
            err = <object> o
            e = type(attr,(Error,),{"err":err, "parent":self})
            self.cache[attr] = e
            return e
        raise AttributeError(attr)
        
    def __getitem__(self, index):
        cdef PyObject *o
        o = PyDict_GetItem(self.reverse, index)
        if o == NULL:
            raise IndexError(index)
        attr = <object> o #this gets us a new reference (release at end)
        return self.__getattr__(attr)
    def preload(self):
        for x in self.mappings:
            self.__getattr__(x)
        
cdef inline error_list_factory(char * name, dict mapping):
    cdef ErrorList l = ErrorList.__new__(ErrorList)
    l.name = name
    l.mappings = mapping
    l.reverse = reverse_dict(mapping)
    l.cache = dict()
    return l

#------------------------------------------------------------
# Implementation of get_info and attribute access
#-------------------------------------------------------------
cdef class Constant:
    cdef ConstantList parent
    cdef value
    
    def __init__(self, *args, **kwargs):
        raise TypeError("Cannot create this type")
    def __str__(self):
        return self.parent.as_string(self.value) 
    def __richcmp__(Constant x, Constant y, int op):
        if op == Py_EQ:
            return x.value == y.value
        elif op == Py_NE:
            return x.value != y.value
        raise TypeError("Only equality supported")
    def __hash__(self):
        return long(self.value)

cdef class IntConstant(Constant):
    def __int__(self):
        return self.value

cdef int_constant_factory(parent, object value):
    cdef IntConstant c =IntConstant.__new__(IntConstant)
    c.parent = parent
    c.value = value
    return c
  
cdef class BitfieldConstant(IntConstant):
    def __or__(BitfieldConstant x, BitfieldConstant y):
        return bitfield_constant_factory(x.parent, x.value | y.value)
    def __and__(BitfieldConstant x, BitfieldConstant y):
        return bitfield_constant_factory(x.parent, x.value & y.value)
    def __xor__(BitfieldConstant x, BitfieldConstant y):
        return bitfield_constant_factory(x.parent, x.value ^ y.value)
    def __invert__(self):
        return bitfield_constant_factory(self.parent, ~self.value)

cdef bitfield_constant_factory( parent, object value):
    cdef BitfieldConstant c = BitfieldConstant.__new__(BitfieldConstant)
    c.parent = parent
    c.value = value
    return c

cdef class ConstantList:
    cdef size_t size
    cdef dict name2value
    cdef dict value2name
    cdef str name
    cdef ConstFactory factory
    
    def __init__(self, *args, **kwargs):
        raise TypeError("Cannot create this type")    
    
    cdef update_names(self):
        if self.value2name is None:
            self.value2name = reverse_dict(self.name2value)

    cdef bint check_value(self, value):
        self.update_names()
        return (value in self.value2name)

    def __getitem__(self, index):
        return self.create_constant(index)
        
    def __str__(self):
        return self.name + ': ' + ','.join(self.name2value.iterkeys())
    def get_size(self):
        return self.size    
    
    cdef  create_from_void(self, void *p):
        return None
    cdef  create_constant(self, value): 
        return  self.factory(self,value)
        
    def __getattr__(self, attr):
        x = self.name2value[attr]
        return self.create_constant(x)

    cdef as_string(self, value):
        self.update_names()
        return self.name + '.' + self.value2name[value]

cdef class IntConstantList(ConstantList):
    def __cinit__(self):
        self.size = sizeof(cl_int)
        self.factory = <ConstFactory> int_constant_factory
    cdef create_from_void(self, void *p):
        return self.create_constant((<cl_int*> p)[0])

cdef IntConstantList int_constant_list_factory(str name, dict name2value):
    cdef IntConstantList l = IntConstantList.__new__(IntConstantList)
    l.name = name
    l.name2value = name2value
    return l

cdef class UIntConstantList(IntConstantList):
    cdef  create_from_void(self, void *p):
        return self.create_constant( (<cl_uint*> p)[0])
    
cdef UIntConstantList uint_constant_list_factory(str name, dict name2value):
    cdef UIntConstantList l = UIntConstantList.__new__(UIntConstantList)
    l.name = name
    l.name2value = name2value    
    return l

cdef class BitfieldList(ConstantList):
    def __cinit__(self):
        self.size = sizeof(cl_bitfield)
        self.factory = <ConstFactory>bitfield_constant_factory
    cdef  create_from_void(self, void *p):
        return self.create_constant( (<cl_bitfield*> p)[0])
        
    cdef get_bits(self, cl_bitfield value):
        cdef cl_bitfield x
        l = []
        for name,x in self.name2value.iteritems():
            if x & value == x:
                l.append(name)
        return l
        
    cdef as_string(self, value):
        cdef cl_bitfield c_value = value
        return self.name + '.' +  "|".join(self.get_bits(c_value))
        
cdef BitfieldList bitfield_list_factory(str name, dict name2value):
    cdef BitfieldList l = BitfieldList.__new__(BitfieldList)
    l.name = name
    l.name2value = name2value
    return l

#------------------------------------------------------------
#creators for different types

cdef create_size_t(void * p, extra) :
    return (<size_t*>  p)[0]

cdef create_bool(void *p, extra):
    return (<cl_bool*> p)[0]
    
cdef create_uint(void *p, extra):
    return (<cl_uint *> p)[0]
    
cdef create_ulong(void *p, extra):
    return (<cl_ulong *> p)[0]

cdef create_int(void *p, extra):
    return (<cl_int *> p) [0]
    
cdef create_getter(void *p, extra):
    cdef type t = extra
    return info_getter_factory(t, (<void_ptr* > p)[0])

cdef create_referenced(void *p,extra):
    cdef type t = extra
    cdef Referenced x = referenced_factory(t, (<void_ptr*>p)[0],None)
    print "Extra retain"
    x.retain() #FIXME: do we need this?
    return x

cdef create_constant(void *p, extra):
    cdef ConstantList clist = extra
    return clist.create_from_void(p)

cdef create_image_format(void *p, extra):
    return image_format_factory(<cl_image_format *>p)


cdef inline object get_value(InfoFunc info, void_ptr id,  
        cl_uint attr, size_t size, CreateFunc creator, extra, void *info_extra):
    cdef u_value result
    cdef cl_int err
    if info_extra == NULL:
         err = info(id, attr, size, &result, NULL)
    else:
        err = (<ExtraInfoFunc> info)(id, info_extra, attr, size, &result, NULL)
    
    check_error("Accessing info", err)
    return creator(&result.v64, extra)
    
cdef inline get_array_value(InfoFunc info, void_ptr id, 
        cl_uint attr, size_t size, CreateFunc creator, extra, void *info_extra):
    cdef size_t return_size
    cdef void * array
    cdef unsigned int x
    cdef cl_int err
    
    if info_extra == NULL:
        err = info(id, attr, 0, NULL, &return_size)
    else:
        err = (<ExtraInfoFunc>info)(id, info_extra, attr, 0, NULL, &return_size)
    check_error("Accessing Info",err)
    array = PyMem_Malloc(return_size)
    
    try:
        if info_extra == NULL:
            err = info(id, attr, return_size, array, NULL)
        else:
            err = (<ExtraInfoFunc>info)(id, info_extra, attr, return_size, array, NULL)
        check_error("Accessing Info", err)
        result = []
        for x from 0 <= x < return_size by size:
            result.append(creator(array+x, extra) )
    finally:
        PyMem_Free(array)    
    return result

#------------------------------------------------------
# Utility classes
#--------------------------------------------------------

#base of most of the API classes.
#Implements attribute access of OpenCL info
cdef class InfoGetter:
    cdef void_ptr id
    cdef InfoFunc info
    cdef readonly ConstantList params
    cdef dict param_types
    
    def __init__(self, *args, **kwargs):
        raise TypeError("Cannot create this type")
    
    def __getattr__(self, attr):
        cdef PyObject * o
        if self.params is  None :
            raise AttributeError(attr)
        o = PyDict_GetItem( self.params.name2value,attr)
        if o == NULL:
            raise AttributeError(attr)
        return self.c_get_info(<object> o, self.info, self.param_types, NULL)

    def get_info(self, Constant param_name):
        return self.c_get_info(param_name.value, self.info, self.param_types, NULL)

    cdef c_get_info(self, cl_uint p_attr, InfoFunc info, dict param_types, void *info_extra):
        cdef cl_uint attr = p_attr
        cdef size_t return_size
        cdef size_t size = 0
        cdef cl_int err
        cdef char* c_str
        cdef u_value value
        cdef ConstantList constant
        cdef InfoGetter item
        cdef CreateFunc creator = NULL
        cdef bint is_array = False
        if param_types is not None and attr in param_types and param_types[attr] != 'char[]':
            t = self.param_types[attr]
            if isinstance(t, tuple):
                is_array = True
                t = t[0]

            if isinstance(t,str):
                if  t[-2:] == '[]':
                    is_array = True
                    t = t[:-2]
                if t == 'cl_bool':
                    creator = create_bool
                    size = sizeof(cl_bool)
                elif t == 'cl_uint':
                    creator = create_uint
                    size = sizeof(cl_uint)
                elif t == 'size_t':
                    creator = create_size_t
                    size = sizeof(size_t)
                elif t == 'cl_ulong':
                    creator = create_ulong
                    size = sizeof(cl_ulong)
                elif t == 'cl_int':
                    creator = create_int
                    size = sizeof(cl_int)
                elif t == 'cl_image_format':
                    creator = create_image_format
                    size = sizeof(cl_image_format)
            elif isinstance(t, ConstantList):
                constant = t
                size = constant.size
                creator = create_constant
            elif isinstance(t,type) and issubclass(t,Referenced):
                size = sizeof(void_ptr)
                creator = create_referenced
            elif isinstance(t,type) and issubclass(t, InfoGetter):
                size = sizeof(void_ptr)
                creator = create_getter
            if creator ==  NULL or size == 0:
                raise TypeError()        
            if is_array:
                return get_array_value(info, self.id, attr, size, creator, t, info_extra)
            return get_value(info, self.id, attr, size, creator, t, info_extra)
        else:
            #string
            if info_extra == NULL:
                err = info(self.id, attr, 0, NULL, &return_size)
            else:
                err = (<ExtraInfoFunc>info)(self.id, info_extra, attr, 0, NULL, &return_size)
            check_error("Accessing Info", err)
            result = PyString_FromStringAndSize(NULL, return_size)
            c_str = PyString_AsString(result)
            
            if info_extra == NULL:
                err = info(self.id, attr, return_size, c_str, NULL)
            else:
                err = (<ExtraInfoFunc>info)(self.id, info_extra, attr, return_size, c_str, NULL)
            check_error("Accessing Info", err)
            return result
            
    cdef size_t c_get(self, cl_uint attr, void *data, size_t size) except? 0:
        cdef size_t ret_size
        check_error("Accessing Info", self.info(self.id, attr, size, data, &ret_size))
        return ret_size   
        
    def __hash__(self):
        return <long>self.id
    def __cmp__(self, InfoGetter other):
        if self.id == other.id:
            return 0
        if self.id < other.id:
            return -1
        return 1

cdef info_getter_factory(type t, void_ptr id):
    cdef InfoGetter i = t.__new__(t)
    i.id = id
    #print "Built:", t, <long long> id
    return i

#Base of all opencl errors.
#Errors are dynamic subclasses of this.
class Error(StandardError):
    err = None
    parent = None
    def __init__(self, value = ''):
        self.value = value
    def __str__(self):
        return "%s( %s)" %(type(self).__name__, str(self.value))
    def __repr__(self):
        cdef ErrorList l = self.parent
        return "%s.%s(%s)" %(l.name, type(self).__name__, repr(self.value))

#Base class for objects that keep a reference count
cdef class Referenced(InfoGetter):
    cdef ReferenceFunc _retain
    cdef ReferenceFunc _release
    cdef cl_uint ref_count_param
    cdef readonly object base
    
    def retain(self):
        check_error( "Retain", self._retain(self.id))
    
    def release(self):
        check_error("Release",  self._release(self.id))

    def __dealloc__(self):
        cdef cl_uint count
        cdef cl_int err
        print "Releasing:", type(self),  PyLong_FromVoidPtr(self.id),
        self.c_get(self.ref_count_param, &count, sizeof(cl_uint))
        print count,
        self._release(self.id)
        print "."
        
cdef referenced_factory(type t, void_ptr id, base):
    cdef Referenced i = info_getter_factory(t,id)
    i.base = base
    #print "Built:", t, <long long> id, base
    return i
    
#Simple class for specifying OpenCL Local Memory sizes
cdef class LocalMemory:
    cdef size_t size
    def __init__(self, size_t size):
        self.size = size
    def __str__(self):
        return "LocalMemory(%d)" % (self.size,)
    def __repr__(self):
        return "<LocalMemory(%d)>" % (self.size,)
    
#base class for memory objects: Buffer Image2D image3d
cdef class MemObject(Referenced):
    cdef readonly mem_base
    cdef Py_buffer view
    types = {
        CL_MEM_TYPE: mem_object_type,
        CL_MEM_FLAGS: mem_flags,
        CL_MEM_SIZE: "size_t",
        #CL_MEM_HOST_PTR: "void *",  #FIXME
        CL_MEM_MAP_COUNT: "cl_uint",
        CL_MEM_REFERENCE_COUNT: "cl_uint",
        CL_MEM_CONTEXT: Context,
    }
    def __cinit__(self):
        self.info = <InfoFunc> clGetMemObjectInfo 
        self._retain = <ReferenceFunc> clRetainMemObject
        self._release = <ReferenceFunc> clReleaseMemObject
        self.ref_count_param = CL_MEM_REFERENCE_COUNT
        self.params = c_mem_info
        self.param_types = self.types
        self.view.buf = NULL
    def __dealloc__(self):
        if self.view.buf != NULL:
            print "Releasing buffer", id(self),
            release_buffer(&self.view)
            print "."
            
cdef mem_object_factory(type t, void_ptr id, base, mem_base, Py_buffer *view):
    cdef MemObject o = referenced_factory(t,id,base)
    o.mem_base = mem_base
    o.view = view[0]
    return o

cdef class Image(MemObject):
    cdef ConstantList image_params
    cdef dict image_param_types
    cdef unsigned int dim
    
    image_types = {
            CL_IMAGE_FORMAT  : 'cl_image_format',
            CL_IMAGE_ELEMENT_SIZE: 'size_t',
            CL_IMAGE_ROW_PITCH : 'size_t',
            CL_IMAGE_SLICE_PITCH : 'size_t',
            CL_IMAGE_WIDTH : 'size_t',
            CL_IMAGE_HEIGHT : 'size_t',
            CL_IMAGE_DEPTH  : 'size_t',
        }
    def __cinit__(self):
        self.image_params = c_image_info
        self.image_param_type = self.image_types
        
    def __getattr__(self, attr):
        cdef PyObject * o
        
        o = PyDict_GetItem(self.image_params.name2value, attr)
        if o == NULL:
            MemObject.__getattr__(self, attr)
        return self.c_get_info(<object> o, <InfoFunc> clGetImageInfo, 
                self.image_param_types, NULL)

    def get_image_info(self, Constant param_name):
        return self.c_get_info(param_name.value, <InfoFunc> clGetImageInfo, self.image_param_types, NULL)        

#implements a buffer interface to mapped memory
cdef class MemMap:
    cdef CommandQueue queue     #for dequeueing and ref counting (though there is a ref in event) 
    cdef MemObject memobj       #for ref-counting
    cdef void * buf             
    cdef Py_ssize_t size
    cdef readonly bint readonly
    cdef readonly bint readable #not enforced. If false, data may be out of date or worse, non-sensical
    cdef readonly Event event   #for waiting on and exposing to client to wait on
    
    def __init__(self):
        raise TypeError("Cannot create type")
    
    def __dealloc__(self):
        self.emergency_unmap()
    cpdef emergency_unmap(self):
        cdef cl_int err
        cdef cl_event event
        cdef Event e
        cdef Constant c
        cdef bytes text
        if self.buf != NULL:
            #PyErr_WarnEx_(NULL, "MemMap not unmapped before deletion. Calling emegency unmap (BAD)", 1)
            PySys_WriteStderr("Warning: MemMap not unmapped before deletion. Calling emegency unmap (BAD)")
            err = clEnqueueUnmapMemObject(self.queue.id,self.memobj.id,  self.buf, 0, NULL, &event)
            self.buf = NULL
            self.size = 0
            try:
                check_error("Emergency unmap", err)
                check_error("Emergency unmap", clWaitForEvents(1, &event))
            finally:
                check_error("Emergency unmap", clReleaseEvent(event))

    def unmap(self):
        if self.buf != NULL:
            return self.queue.enqueue_unmap_mem_object(self) 
        raise TypeError("Buffer already unmapped")
        
    def status(self):
        cdef cl_int c
        self.event.c_get(CL_EVENT_COMMAND_EXECUTION_STATUS, &c, sizeof(cl_int))
        return c_command_status.create_constant(c)
    def ready(self):
        cdef cl_int c
        self.event.c_get(CL_EVENT_COMMAND_EXECUTION_STATUS, &c, sizeof(cl_int))
        return c == CL_COMPLETE
    def wait(self):
        self.event.wait()
    #python 2 buffer interface
    def __getreadbuffer__(self, Py_ssize_t seg, void **ptr):
        ptr[0] = self.buf
        return self.size
    def __getwritebuffer__(self, Py_ssize_t seg, void **ptr):
        if self.readonly:
            raise TypeError("Memory readonly")
        ptr[0] =self.buf
        return self.size
    def __getcharbuffer__(self, Py_ssize_t seg, char **ptr):
        ptr[0] = <char *> self.buf
        return self.size
    def __getsegcount__(self, int *lenp):
        if lenp != NULL:
            lenp[0] = self.size
        if self.buf == NULL:
            return 0
        return 1
    # new buffer inteface (python 2.6 and up)
    def __getbuffer__(self, Py_buffer *view, int flags):
        PyBuffer_FillInfo(view, self, self.buf, self.size, self.readonly, flags) 
        
    def __repr__(self):
        return "<MemMap: %d (%d)>" %(PyLong_FromVoidPtr(self.buf), self.size)
        
cdef MemMap mem_map_factory(CommandQueue queue, MemObject memobj, cl_event c_event, void* buf, size_t size, cl_bitfield flags):
    cdef Event event = referenced_factory(Event, c_event, queue)
    cdef MemMap o = MemMap.__new__(MemMap)
    o.queue = queue
    o.memobj = memobj
    o.event = event
    o.buf  = buf
    o.size = size
    o.readonly = not flags | CL_MAP_WRITE
    o.readable = flags | CL_MAP_READ
    return o
 
cdef class StridedMemMap(MemMap):
    cdef int dim
    cdef Py_ssize_t strides[MAX_DIM+1]
    cdef Py_ssize_t shape[MAX_DIM+1]
    cdef Py_ssize_t item_size
    cdef char format[6]
    cdef char fort
    cdef bint contig
    
    #python 2 buffer interface
    cdef Py_ssize_t c_getreadbuffer(self, Py_ssize_t seg, void **ptr):
        cdef Py_ssize_t row
        cdef Py_ssize_t slice
        cdef int dim = self.dim
        
        if self.fort == 'C':
            row = seg % self.shape[dim - 1] 
            slice = seg / self.shape[dim-1]
            ptr[0] = self.buf + row * self.strides[dim-1] + slice * self.strides[dim-2]
            return self.shape[dim] * self.item_size
        else:
            row = seg % self.shape[2]  # height
            slice = seg / self.shape[2] #height
        
            ptr[0] = self.buf + row * self.strides[2] + slice * self.strides[3]
            return self.shape[1] * self.item_size         
    def __getreadbuffer__(self, Py_ssize_t seg, void **ptr):
        return self.c_getreadbuffer(seg,ptr)
    def __getwritebuffer__(self, Py_ssize_t seg, void **ptr):
        if self.readonly:
            raise TypeError("Memory readonly")
        return self.c_getreadbuffer( seg, ptr)
    def __getcharbuffer__(self, Py_ssize_t seg, char **ptr):
        return self.c_getreadbuffer(seg,<void **> ptr)
    def __getsegcount__(self, int *lenp):
        if lenp != NULL:
            lenp[0] = self.size
        
        #depth is 1 if dim == 2
        if self.fort == 'C':
            return self.shape[self.dim -1 ] * self.shape[self.dim-2] 
        else:
            return self.shape[2] * self.shape[3] 
    # new buffer inteface (python 2.6 and up)
    def __getbuffer__(self, Py_buffer *view, int flags) :
        if view == NULL:
            return

        if ((flags & PyBUF_WRITABLE) == PyBUF_WRITABLE) and self.readonly :
            raise BufferError("Object is not writable.")
        view.obj = self            
        view.buf = self.buf
        view.len = self.size
        view.readonly = self.readonly
        view.itemsize = self.item_size
        view.format = NULL
        if ((flags & PyBUF_FORMAT) == PyBUF_FORMAT):
            view.format = self.format
        view.ndim  = self.dim
        view.shape = NULL
        if ((flags & PyBUF_ND) == PyBUF_ND):
            view.shape =  (self.shape + 1);
        view.strides = NULL
        if ((flags & PyBUF_STRIDES) == PyBUF_STRIDES):
            view.strides = (self.strides + 1);
        view.suboffsets = NULL;
        view.internal = NULL;
	    
	    
        if (((flags & PyBUF_C_CONTIGUOUS) == PyBUF_C_CONTIGUOUS) or 
                ((flags & PyBUF_STRIDES) != PyBUF_STRIDES)) and  (self.contig == False or self.fort != 'C'):
            raise BufferError("Object is not C contiguous") 
        elif ((flags & PyBUF_F_CONTIGUOUS) == PyBUF_F_CONTIGUOUS) and (self.config == False or self.fort == 'C'):
            raise BufferError("Object is not F contiguous")
        elif ((flags & PyBUF_ANY_CONTIGUOUS) == PyBUF_ANY_CONTIGUOUS) and self.contig == False:
            raise BufferError("Object is not contiguous")


	    
cdef StridedMemMap strided_mem_map_factory(CommandQueue queue, MemObject memobj, cl_event c_event, void* buf, size_t size, 
            int dim, cl_image_format *format, size_t row_pitch, size_t slice_pitch, size_t *region, char fort, cl_bitfield flags):
    cdef Event event = referenced_factory(Event, c_event, queue)
    cdef StridedMemMap o = StridedMemMap.__new__(StridedMemMap)
    cdef size_t t
    o.queue = queue
    o.memobj = memobj
    o.event = event
    o.buf  = buf
    o.size = size
    o.readonly = not flags | CL_MAP_WRITE
    o.readable = flags | CL_MAP_READ
    o.dim = dim
    o.item_size = calc_image_format(format, o.format)
    o.fort = fort   
    
    #check contiguity
    o.contig = False
    t = region[0] * o.item_size    
    if row_pitch == t and (dim == 2 or slice_pitch == t*region[1]):
        o.contig = True
        
    #setup shape and strides
    if fort == 'C':
        o.strides[dim] = o.item_size
        o.strides[dim - 1] = row_pitch
        o.strides[dim - 2] = slice_pitch
        
        o.shape[dim] = region[0] #width
        o.shape[dim-1] = region[1] #height
        o.shape[dim-2] = region[2]  #depth
    else:
        o.shape[1] = region[0]
        o.shape[2] = region[1]
        o.shape[3] = region[2] #not used if dim == 2
    
        o.strides[1] = o.item_size
        o.strides[2] = row_pitch        
        o.strides[3] = slice_pitch # not used if dim == 2
    return o
 
cdef class ImageFormat:
    cdef cl_image_format format
    cdef readonly Constant CHANNEL_ORDER
    cdef readonly Constant CHANNEL_DATA_TYPE
    def __init__(self, Constant order, Constant data_type):
        self.format.image_channel_order = order.value
        format.image_channel_data_type = data_type.value
        self.CHANNEL_ORDER = order
        self.CHANNEL_DATA_TYPE = data_type
    
cdef image_format_factory(cl_image_format *f):
    cdef ImageFormat format = ImageFormat.__new__(ImageFormat)
    format.format = f[0]
    return format
 
#-------------------------------------------------
# API Class definitions
#---------------------------------------------
cdef class Platform(InfoGetter):
    def __cinit__(self):
        self.info = <InfoFunc> clGetPlatformInfo
        self.params = c_platform_info
        self.param_types = None  #all strings
        
    def get_devices(self, cl_uint device_type = CL_DEVICE_TYPE_ALL):
        cdef cl_int err
        cdef cl_uint return_size
        cdef cl_device_id *ids
        cdef int x
        cdef Device d
        
        check_error("Getting Devices", clGetDeviceIDs(self.id, device_type, 0, NULL, &return_size))
        ids = <cl_device_id *> PyMem_Malloc(sizeof(cl_device_id) * return_size)
        try:
            check_error("Getting Devices", clGetDeviceIDs(self.id,device_type, return_size, ids, NULL))
            result = []
            for x in range(return_size):
                d = info_getter_factory(Device, ids[x])
                result.append(d)
        finally:
            PyMem_Free(ids)
        return result    
    def get_device_ids(self, cl_uint device_type = CL_DEVICE_TYPE_ALL):
        return self.get_devices(device_type)    
    def create_context(self, properties = None, devices = None):
        return create_context( properties, devices)
    
    def create_context_from_type(self, properties=None, device_type = CL_DEVICE_TYPE_DEFAULT):
        return create_context_from_type(properties, device_type, self)


cdef class Device(InfoGetter):
    types = {
            CL_DEVICE_TYPE: device_type,
            CL_DEVICE_VENDOR_ID: 'cl_uint',
            CL_DEVICE_MAX_COMPUTE_UNITS: 'cl_uint',
            CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS: 'cl_uint',
            CL_DEVICE_MAX_WORK_ITEM_SIZES: 'size_t[]',
            CL_DEVICE_MAX_WORK_GROUP_SIZE: 'size_t',
            CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR: 'cl_uint',
            CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT: 'cl_uint',
            CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT: 'cl_uint',
            CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG: 'cl_uint',
            CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT: 'cl_uint',
            CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE: 'cl_uint',
            CL_DEVICE_MAX_CLOCK_FREQUENCY: 'cl_uint',
            CL_DEVICE_ADDRESS_BITS: 'cl_uint',
            CL_DEVICE_MAX_MEM_ALLOC_SIZE: 'cl_ulong',
            CL_DEVICE_IMAGE_SUPPORT: 'cl_bool',
            CL_DEVICE_MAX_READ_IMAGE_ARGS: 'cl_uint',
            CL_DEVICE_MAX_WRITE_IMAGE_ARGS: 'cl_uint',
            CL_DEVICE_IMAGE2D_MAX_WIDTH: 'size_t',
            CL_DEVICE_IMAGE2D_MAX_HEIGHT: 'size_t',
            CL_DEVICE_IMAGE3D_MAX_WIDTH: 'size_t',
            CL_DEVICE_IMAGE3D_MAX_DEPTH: 'size_t',
            CL_DEVICE_MAX_SAMPLERS: 'cl_uint',
            CL_DEVICE_MAX_PARAMETER_SIZE: 'size_t',
            CL_DEVICE_MEM_BASE_ADDR_ALIGN: 'cl_uint',
            CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE: 'cl_uint',
            CL_DEVICE_SINGLE_FP_CONFIG: c_device_fp_config,
            CL_DEVICE_GLOBAL_MEM_CACHE_TYPE: c_device_mem_cache_type,
            CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE: 'cl_uint',
            CL_DEVICE_GLOBAL_MEM_CACHE_SIZE: 'cl_ulong',
            CL_DEVICE_GLOBAL_MEM_SIZE: 'cl_ulong',
            CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE: 'cl_ulong',
            CL_DEVICE_MAX_CONSTANT_ARGS: 'cl_uint',
            CL_DEVICE_LOCAL_MEM_TYPE: c_device_local_mem_type,
            CL_DEVICE_LOCAL_MEM_SIZE: 'cl_ulong',
            CL_DEVICE_ERROR_CORRECTION_SUPPORT: 'cl_bool',
            CL_DEVICE_PROFILING_TIMER_RESOLUTION: 'size_t',
            CL_DEVICE_ENDIAN_LITTLE: 'cl_bool',
            CL_DEVICE_AVAILABLE: 'cl_bool',
            CL_DEVICE_COMPILER_AVAILABLE: 'cl_bool',
            CL_DEVICE_EXECUTION_CAPABILITIES: c_device_exec_capabilities,
            CL_DEVICE_QUEUE_PROPERTIES: c_command_queue_properties,
            CL_DEVICE_PLATFORM: Platform,
        }
    def __cinit__(self):
        self.info = <InfoFunc> clGetDeviceInfo
        self.params = c_device_info
        self.param_types = self.types
        
    def create_context(self, properties = None):
        return create_context(properties, self)

cdef class Context(Referenced):
    types = {
            CL_CONTEXT_REFERENCE_COUNT: 'cl_uint',
            CL_CONTEXT_DEVICES: (Device,),
            CL_CONTEXT_PROPERTIES: (c_context_properties,),
        }
    def __cinit__(self):
        self.info = <InfoFunc> clGetContextInfo
        self._retain = <ReferenceFunc> clRetainContext
        self._release = <ReferenceFunc> clReleaseContext
        self.ref_count_param = CL_CONTEXT_REFERENCE_COUNT
        self.params = c_context_info
        self.param_types = self.types

    def create_command_queue(self, device,  properties = None):
        cdef cl_command_queue_properties c_props
        cdef cl_command_queue queue
        cdef void_ptr device_id
        cdef cl_int err
        if properties is None:
            c_props = 0
        else:
            c_props = properties
        if device is None:
            raise ValueError("Device cannot be none")
        device_id = (<Device> device).id
        
        queue = clCreateCommandQueue(self.id, device_id, c_props, &err)
        check_error("Create Command Queue", err)
        return referenced_factory(CommandQueue,  queue, self)
    
    #if data is a buffer and CL_USE_HOST_PTR is specified or infered, the pythoncl buffer keeps a reference
    #   to the input buffer and view.  It releases these on deletion.
    def create_buffer(self,  BitfieldConstant flags=None, size_t size = 0, host_ptr=None):
        cdef cl_bitfield c_flags = CL_MEM_READ_WRITE
        cdef cl_bitfield WRITABLE = (CL_MEM_READ_WRITE | CL_MEM_WRITE_ONLY)
        cdef void * c_ptr = NULL
        cdef Py_buffer view
        cdef Buffer result
        cdef cl_mem buf
        cdef cl_int err
        cdef object mem_base = None
        cdef Py_ssize_t l        
        
        view.buf = NULL
        
        if flags is not None:
            c_flags =  flags.value
        if host_ptr is None:
            if size == 0:
                raise ValueError("Must specify size if data is None")
            if flags is None:
                c_flags |= CL_MEM_ALLOC_HOST_PTR
        else:
            l = get_buffer(host_ptr, (c_flags & WRITABLE) and (c_flags & CL_MEM_USE_HOST_PTR),
                            &c_ptr, &view)
            mem_base = host_ptr
            if size != 0:
                if size > l:
                    release_buffer(&view)
                    raise ValueError("Specified size is greater than data length")
            else:
                size = l
            if flags is None:
                if view.buf == NULL or view.readonly:
                    #default flags are read/write so we need to copy
                    c_flags |= (CL_MEM_COPY_HOST_PTR | CL_MEM_ALLOC_HOST_PTR)
                else:
                    #no need to copy. This *might* skip a copy, but is also slower.
                    c_flags |= CL_MEM_USE_HOST_PTR
            
        #print c_flags, c_size, PyLong_FromVoidPtr(host_ptr)
        with nogil:
            buf = clCreateBuffer(self.id, c_flags, size, c_ptr,&err)
        
        if (err != CL_SUCCESS):
            release_buffer(&view)
            raise_error("Creating Buffer", err)
        
        if (mem_base is not None) and (c_flags & CL_MEM_COPY_HOST_PTR):
            #data copied. Release buffer
            release_buffer(&view)
            mem_base = None
        try:        
            result = mem_object_factory(Buffer, buf, self, mem_base, &view)
        except:
            release_buffer(&view)
            raise
        return result 
    def create_program_with_sources(self, strings):
        cdef cl_int err
        cdef char ** c_sources
        cdef cl_uint count
        cdef size_t *lengths
        cdef unsigned int i
        cdef cl_program prog
        cdef Py_ssize_t s
        
        if isinstance (strings, str):
            strings = tuple(strings)
            
        count = len(strings)
        
        try:
            c_sources =  <char **>PyMem_Malloc(sizeof(char *) * count)
            lengths = <size_t *> PyMem_Malloc(sizeof(size_t) * count)    
        
            for i,source in enumerate(strings):
                PyString_AsStringAndSize(source, <char **> &(c_sources[i]), &s)
                lengths[i] = s
                
            prog = clCreateProgramWithSource(self.id, count, c_sources, lengths, &err)
            check_error("Creating Program", err)
        finally:
            PyMem_Free(c_sources)
            PyMem_Free(lengths)
        
        return referenced_factory(Program, prog, self)
    def get_supported_image_formats(self, BitfieldConstant flags=None, Constant image_type = None):
        cdef cl_image_format *formats
        cdef cl_uint num_formats
        cdef cl_int err
        cdef cl_bitfield c_flags = 0
        cdef int i
        
        if image_type is None:
            raise ValueError("Must specify image_type")
        
        if flags is not None:
            c_flags = flags.value
        
        check_error('Get Supported Image Formats',
            clGetSupportedImageFormats(self.id, c_flags, image_type.values, 0, NULL, &num_formats))
        
        formats = <cl_image_format *> PyMem_Malloc(sizeof(cl_image_format) * num_formats)
        
        try:
            check_error('Get Supported Image Formats',
                clGetSupportedImageFormats(self.id, flags.value, image_type.values, num_formats, formats, NULL))
        
            results = []
            for i in range(num_formats):
                results.append(image_format_factory(formats +  i))
        finally:
            PyMem_Free(formats)
        return results
    
    def create_image_2d(self, BitfieldConstant flags = None, ImageFormat image_format = None, 
            size_t image_width = 0, size_t image_height= 0, size_t image_row_pitch = 0,  host_ptr = None):
        cdef cl_bitfield c_flags = CL_MEM_READ_WRITE
        cdef cl_bitfield WRITABLE = (CL_MEM_READ_WRITE | CL_MEM_WRITE_ONLY)
        cdef void * c_ptr = NULL
        cdef Py_buffer view
        cdef Image2D result
        cdef cl_mem buf
        cdef cl_int err
        cdef object mem_base = None
        cdef Py_ssize_t l        
        
        if image_format is None:
            raise ValueError("Image must not be None")
        
        view.buf = NULL
        
        if (image_width == 0 or image_height ==0):
            raise ValueError("Width and height must be nonzero")
        if flags is not None:
            c_flags =  flags.value
        if host_ptr is None:
            if flags is None:
                c_flags |= CL_MEM_ALLOC_HOST_PTR
        else:
            l = get_buffer(host_ptr, (c_flags & WRITABLE) and (c_flags & CL_MEM_USE_HOST_PTR),
                            &c_ptr, &view)
                                            
            mem_base = host_ptr
            if (image_width * image_height) > l:
                release_buffer(&view)
                raise ValueError("Specified size is greater than data length")
            
            if flags is None:
                if view.buf == NULL or view.readonly:
                    #If view not available, assume read-only
                    #because the default flags are read/write, we need opencl to make a copy of readonly buffer
                    c_flags |= (CL_MEM_COPY_HOST_PTR | CL_MEM_ALLOC_HOST_PTR)
                else:
                    #no need to copy. This *might* skip a copy, but is also slower.
                    c_flags |= CL_MEM_USE_HOST_PTR
            
        #print c_flags, c_size, PyLong_FromVoidPtr(host_ptr)
        with nogil:
            buf = clCreateImage2D(self.id, c_flags, &image_format.format, image_width, image_height, image_row_pitch, c_ptr,&err)
        
        if (err != CL_SUCCESS):
            release_buffer(&view)
            raise_error("Creating Buffer", err)
        
        if (mem_base is not None) and (c_flags & CL_MEM_COPY_HOST_PTR):
            #data copied. Release buffer
            release_buffer(&view)
            mem_base = None
        try:        
            result = mem_object_factory(Image2D, buf, self, mem_base, &view)
        except:
            release_buffer(&view)
            raise
        return result 
    def create_image_3d(self, BitfieldConstant flags = None, ImageFormat format = None, 
            size_t image_width = 0, size_t image_height = 0, size_t image_depth = 0, 
            size_t image_row_pitch = 0,size_t image_slice_pitch = 0, host_ptr = None):
        cdef cl_bitfield c_flags = CL_MEM_READ_WRITE
        cdef cl_bitfield WRITABLE = (CL_MEM_READ_WRITE | CL_MEM_WRITE_ONLY)
        cdef void * c_ptr = NULL
        cdef Py_buffer view
        cdef Image3D result
        cdef cl_mem buf
        cdef cl_int err
        cdef object mem_base = None
        cdef Py_ssize_t l        
        
        if format is None:
            raise ValueError("Format cannot be none")
        
        view.buf = NULL
        
        if (image_width == 0 or image_height ==0 or image_depth== 0):
            raise ValueError("Width, height, and depth must be nonzero")
        if flags is not None:
            c_flags =  flags.value
        if host_ptr is None:
            if flags is None:
                c_flags |= CL_MEM_ALLOC_HOST_PTR
        else:
            l = get_buffer(host_ptr, (c_flags & WRITABLE) and (c_flags & CL_MEM_USE_HOST_PTR),
                            &c_ptr, &view)
                                            
            mem_base = host_ptr
            if (image_width * image_height * image_depth) > l:
                release_buffer(&view)
                raise ValueError("Specified size is greater than data length")
            
            if flags is None:
                if view.buf == NULL or view.readonly:
                    #If view not available, assume read-only
                    #because the default flags are read/write, we need opencl to make a copy of readonly buffer
                    c_flags |= (CL_MEM_COPY_HOST_PTR | CL_MEM_ALLOC_HOST_PTR)
                else:
                    #no need to copy. This *might* skip a copy, but is also slower.
                    c_flags |= CL_MEM_USE_HOST_PTR
            
        #print c_flags, c_size, PyLong_FromVoidPtr(host_ptr)
        with nogil:
            buf = clCreateImage3D(self.id, c_flags, &format.format, image_width, image_height, image_depth, 
                                    image_row_pitch, image_slice_pitch, c_ptr,&err)
        if (err != CL_SUCCESS):
            release_buffer(&view)
            raise_error("Creating Buffer", err)
        
        if (mem_base is not None) and (c_flags & CL_MEM_COPY_HOST_PTR):
            #data copied. Release buffer
            release_buffer(&view)
            mem_base = None
        try:        
            result = mem_object_factory(Image3D, buf, self, mem_base, &view)
        except:
            release_buffer(&view)
            raise
        return result 
    def create_sampler(self, bint normalized_coords, Constant addressing_mode, Constant filter_mode):
        cdef cl_int err
        cdef cl_sampler sampler
        
        sampler = clCreateSampler(self.id, normalized_coords, addressing_mode.value, filter_mode.value, &err)
        
        check_error("Create sampler", err)
        
        return referenced_factory(Sampler, sampler, self)
        
                                
cdef class CommandQueue(Referenced):
    types = {
        CL_QUEUE_CONTEXT: Context,
        CL_QUEUE_DEVICE: Device,
        CL_QUEUE_REFERENCE_COUNT: 'cl_uint',
        CL_QUEUE_PROPERTIES: c_command_queue_properties,
    }
    def __cinit__(self):
        self.info = <InfoFunc> clGetCommandQueueInfo
        self._retain = <ReferenceFunc> clRetainCommandQueue
        self._release = <ReferenceFunc> clReleaseCommandQueue
        self.ref_count_param = CL_QUEUE_REFERENCE_COUNT
        self.params = c_command_queue_info
        self.param_types = self.types
    def __setattr__(self, attr, value):
        global c_command_queue_properties
        cdef ConstantList l = c_command_queue_properties
        cdef BitfieldConstant prop
        if attr not in l.name2value:
            raise AttributeError(attr)
        prop = l.create_constant(l.name2value[attr])
        self.set_property( prop, value)
            
    cpdef set_property(self, Constant properties, bint enable):
        cdef cl_int err
        cdef cl_bitfield c_result
        global c_command_queue_properties
        cdef ConstantList l = c_command_queue_properties
        
        with nogil:
            #blocks until all queued commands finish
            err = clSetCommandQueueProperty(self.id, properties.value, enable, &c_result) 
        check_error("Set Property", err)
        
        return l.create_constant(c_result)
        
    #If this call is not blocking, the memory from the input buffer object must not be touched or deallocated
    # before the command actually executes.  It is up to the caller to ensure this. This function releases the 
    # buffer object before returning.    
    def enqueue_read_buffer(self, Buffer buffer, bint blocking_read = True,  
             size_t offset = 0, size_t cb = 0, ptr = None, event_wait_list = None):
        cdef Py_buffer view   
        cdef cl_event c_event
        cdef void *c_ptr
        cdef Py_ssize_t l

        if ptr is None:
            raise ValueError("ptr cannot be None")

        if cb == 0:
            buffer.c_get(CL_MEM_SIZE, &cb, sizeof(size_t))

        l = get_buffer(ptr, True, &c_ptr, &view)

        try:
            if cb > 0 and (cb +offset)> l:
                raise ValueError("Cannot read more data than size of array")
        
            c_event = self.read_buffer(buffer.id, c_ptr, offset, cb, blocking_read,event_wait_list)
        finally:
            release_buffer(&view)
        return referenced_factory(Event, c_event, self)

    #allocate a string and read the buffer into it.
    def enqueue_read_buffer_as_str(self, Buffer buffer, size_t cb=0, event_wait_list = None):
        cdef cl_event event
        if cb == 0:
            buffer.c_get(CL_MEM_SIZE, &cb, sizeof(size_t))
    
        cdef object s = PyString_FromStringAndSize(NULL, cb)
        cdef void *data = PyString_AsString(s)
        #possible problem here:
        #   data is pointer to a python string. It has not been given out yet so we can modify it (per C-API spec)
        #   However, read_buffer releases the Global Interpreter Lock. This may be very bad as we are modifying
        #   the memory of a python object. What if the data is moved by the memory manager?
        event = self.read_buffer(buffer.id, data, 0, cb, True, event_wait_list)
        check_error("Releasing Event", clReleaseEvent(event))  
        return s

    #helper to do heavy lifting of reading a buffer
    cdef cl_event read_buffer(self, cl_mem buf, void *data, size_t offset, size_t size, bint blocking, wait_events) except NULL:
        cdef cl_event event
        cdef cl_event *c_wait = NULL
        cdef cl_uint num_events =0 
        cdef cl_int err
        
        num_events = create_events(wait_events, &c_wait)
        try:
            with nogil:
                err = clEnqueueReadBuffer(self.id, buf, blocking, 
                    offset, size, data, num_events, c_wait, &event)
            check_error("Enqueuing Read Buffer", err)
        finally:
            PyMem_Free(c_wait)
        return event 
    
    
    #If this call is not blocking, the memory from the input buffer object must not be touched or deallocated
    # before the command actually executes.  It is up to the caller to ensure this. This function releases the 
    # buffer object before returning.  
    def enqueue_write_buffer(self, Buffer buffer, bint blocking_write = True, 
            size_t offset = 0, size_t cb = 0, ptr = None , event_wait_list = None):
        cdef void *host_ptr  = NULL
        cdef Py_ssize_t l
        cdef cl_event event
        cdef cl_int err
        cdef cl_event * c_wait = NULL
        cdef cl_uint num_events
        cdef Py_buffer view
        
        if ptr is None:
            raise ValueError("ptr cannot be None")
        
        l = get_buffer(ptr,False, &host_ptr, &view)

        try:
            if cb != 0: 
                if (cb + offset) < l:
                    raise ValueError("Size speficied is greater than size of data")
            else:
                cb = l
            
            num_events=create_events(event_wait_list, &c_wait)
            
            with nogil:
                err = clEnqueueWriteBuffer(self.id, buffer.id, blocking_write, offset, cb, host_ptr, num_events, c_wait, &event)
            check_error("Enqueuing Write Buffer", err)
        finally:
            release_buffer(&view)
            PyMem_Free(c_wait)
        
        return referenced_factory(Event, event, self)

    def enqueue_copy_buffer(self, Buffer src_buffer, Buffer dest_buffer, size_t src_offset = 0, size_t dest_offset = 0,
                            size_t cb = 0, event_wait_list=None):
        cdef cl_event * c_wait
        cdef cl_uint num_events
        cdef cl_event event
        
        num_events = create_events(event_wait_list, &c_wait)
        
        try:
            check_error("Enqueue Copy Buffer",
                clEnqueueCopyBuffer(self.id, src_buffer.id, dest_buffer.id, src_offset, dest_offset, cb, 
                    num_events, c_wait, &event))
        finally:
            PyMem_Free(c_wait)
        
        return referenced_factory(Event, event, self)
        
    def enqueue_read_image(self, Image image, bint blocking_read = True , origin = None, region = None, 
            size_t row_pitch = 0, size_t slice_pitch = 0, ptr = None, event_wait_list = None):
        cdef size_t c_origin[3]
        cdef size_t c_region[3]
        cdef int i
        cdef Py_ssize_t l
        cdef Py_buffer view
        cdef void * c_ptr
        cdef cl_event c_event
        
        if ptr is None:
            raise ValueError("ptr cannot be None")
        
        get_image_params(image.dim, image.id, origin, region, NULL, c_origin, c_region)
        l = get_buffer(ptr, True, &c_ptr, &view)
        
        try:
            if ((c_origin[0] + c_region[0]) *
                (c_origin[1] + c_region[1]) *
                (c_origin[2] + c_region[2])) > l:
                raise ValueError("Cannot read more data than size of array")
        
            c_event = self.read_image(image, c_ptr, c_origin, c_region, row_pitch, slice_pitch, 
                blocking_read,event_wait_list)
        finally:
            release_buffer(&view)
        return referenced_factory(Event, c_event, self)
    
    cdef cl_event read_image(self, Image image, void *data, size_t *origin, size_t *region, 
            size_t row_pitch, size_t slice_pitch, bint blocking, wait_events):
        cdef cl_int err
        cdef cl_event *c_wait
        cdef cl_event event
        cdef cl_uint num_events
        
        num_events = create_events(wait_events, &c_wait)
        
        try:
            with nogil:
                err = clEnqueueReadImage(self.id, image.id, blocking, origin, region, row_pitch, slice_pitch, data, 
                    num_events, c_wait, &event)
        finally:
            PyMem_Free(c_wait)
        
        check_error("Enqueue Read image", err)
        
        return event
        
    def enqueue_write_image(self, Image image,bint blocking_write = True , origin = None, region = None,
            size_t input_row_pitch = 0, size_t input_slice_pitch = 0, ptr = None, event_wait_list = None):
        cdef size_t c_origin[3]
        cdef size_t c_region[3] 
        cdef Py_buffer view
        cdef Py_buffer *default_view = NULL
        cdef void * c_ptr
        cdef Py_ssize_t l
        cdef cl_event *c_wait
        cdef cl_uint num_events
        cdef cl_event event
        cdef cl_int err
        
        if (ptr is None):
            raise ValueError("Ptr cannot be None")
        
        l = get_buffer(ptr, False, &c_ptr, &view, PyBUF_ND)
        
        try:
            if region is None:
                if view.buf !=  NULL and view.ndim == image.dim:
                    default_view = &view
                else:
                    raise ValueError("Cannot determine region size for buffer")     
            
            get_image_params(image.dim, image.id, origin,  region, default_view,c_origin, c_region)
                
            if ((c_origin[0] + c_region[0]) *
                (c_origin[1] + c_region[1]) *
                (c_origin[2] + c_region[2])) > l:
                raise ValueError("Cannot read more data than size of array")
            
            num_events=create_events(event_wait_list, &c_wait)
            
            with nogil:
                err = clEnqueueWriteImage(self.id, image.id, blocking_write, c_origin, c_region,
                        input_row_pitch, input_slice_pitch, c_ptr, num_events, c_wait, &event)
        finally:
            release_buffer(&view)
            PyMem_Free(c_wait)
        check_error("Enqueuing Write Buffer", err)
        return referenced_factory(Event, event, self)

    def equeue_copy_image(self):
        pass
                
    def enqueue_copy_buffer_to_image(self):
        pass
    def enqueue_copy_image_to_buffer(self):
        pass    

    def enqueue_map_buffer(self, Buffer buffer, bint blocking_map = True, map_flags = None,
            size_t offset = 0, size_t cb=0,  event_wait_list = None):
        cdef cl_event *c_wait = NULL
        cdef cl_uint num_events
        cdef cl_bitfield c_flags = CL_MAP_READ | CL_MAP_WRITE
        cdef cl_event c_event
        cdef void *map_buf
        cdef cl_int err
        cdef Event event
        
        if cb == 0:
            buffer.c_get(CL_MEM_SIZE, &cb, sizeof(size_t))
        
        if map_flags is not None:
            c_flags = map_flags
        
        num_events = create_events(event_wait_list, &c_wait)

        try:
            with nogil:
                map_buf = clEnqueueMapBuffer(self.id,buffer.id, blocking_map, c_flags, offset, cb, 
                    num_events, c_wait, & c_event, &err)
            
            check_error("Enqueue map buffer", err)
        finally:
            PyMem_Free(c_wait)
        
        return mem_map_factory(self, buffer, c_event,  map_buf, cb, c_flags)
        
    def enqueue_map_image(self):
#        cdef cl_event *c_wait = NULL
#        cdef cl_uint num_events
#        cdef cl_bitfield c_flags = CL_MAP_READ | CL_MAP_WRITE
#        cdef cl_event c_event
#        cdef void *map_buf
#        cdef cl_int err
#        cdef Event event
#        
#        if flags is not None:
#            c_flags = flags
#        
#        num_events = create_events(wait_events, &c_wait)

#        try:
#            with nogil:
#                map_buf = clEnqueueMapBuffer(self.id,buf.id, blocking, c_flags, offset, size, 
#                    num_events, c_wait, & c_event, &err)
#            
#            check_error("Enqueue map buffer", err)
#        finally:
#            PyMem_Free(c_wait)
#        
#        return mem_map_factory(self, buf, c_event,  map_buf, size, c_flags)
        pass
            
    def enqueue_unmap_mem_object(self, MemMap memobj, event_wait_list = None):
        cdef cl_event * c_wait = NULL
        cdef cl_uint num_events
        cdef cl_int err
        cdef cl_event event
        
        num_events = create_events(event_wait_list, &c_wait)

        try:
            with nogil:
                err = clEnqueueUnmapMemObject(self.id, memobj.memobj.id, memobj.buf, num_events, c_wait, &event)
        finally:
            PyMem_Free(c_wait)
        check_error("Enqueue unmap memory", err)
        memobj.buf = NULL
        memobj.size = 0
        return referenced_factory(Event, event, self)

    def enqueue_nd_range_kernel(self, Kernel kernel, size_t work_dim,global_work_offset = None, 
            global_work_size = None, local_work_size = None, event_wait_list = None):
        cdef cl_event * c_wait
        cdef cl_uint num_events
        cdef size_t size
        cdef int i
        cdef size_t global_size[MAX_DIM]
        cdef size_t local_size[MAX_DIM]
        cdef size_t global_offset[MAX_DIM]
        cdef size_t *global_size_ptr = global_size
        cdef size_t *local_size_ptr = local_size
        cdef size_t *global_offset_ptr = global_offset
        cdef cl_event event
        cdef cl_int err
        
        if work_dim > MAX_DIM:
            raise ValueError("Maximum Dimension %d exceeded: %d" %(MAX_DIM, work_dim))
        
        if global_work_size is None:
            raise ValueError("Global work size cannot be None")
        
        if global_work_offset is None:
            global_offset_ptr = NULL
        else:
            for i in range(work_dim):
                global_offset[i] = global_work_offset[i]
        if local_work_size is None:
            local_size_ptr = NULL
        else:
            for i in range(work_dim):

                local_size[i] =local_work_size[i]
        for i in range(work_dim):
            global_size[i] = global_work_size[i]     
        
        num_events = create_events(event_wait_list, &c_wait)
        
        try:
            with nogil:     
                #does this block?
                err = clEnqueueNDRangeKernel(self.id, kernel.id, work_dim,
                    global_offset_ptr, global_size_ptr, local_size_ptr,
                    num_events, c_wait, &event)
        
        finally:
            PyMem_Free(c_wait)
        check_error("Enqueuing NDRange Kernel", err)
    
        return referenced_factory(Event, event, self)

    def enqueue_task(self, Kernel kernel, event_wait_list=None):
        cdef cl_uint num_events
        cdef cl_event * c_wait = NULL
        cdef cl_int err
        cdef cl_event event
        
        num_events = create_events(event_wait_list, &c_wait)
        
        try:
            with nogil:
                err = clEnqueueTask(self.id, kernel.id, num_events, c_wait, &event)
        finally:
            PyMem_Free(c_wait)
        
        check_error("Enqueueing task", err)
        
        return referenced_factory(Event, event, self)
        
    def enqueue_marker(self):
        cdef cl_event event
        cdef cl_int err
        
        with nogil:
            err = clEnqueueMarker(self.id, &event)
            
        check_error("Enqueueing Marker", err)
        
        return referenced_factory(Event, event, self)
        
    def enqueue_wait_for_events(self,event_list):
        cdef cl_uint num_events
        cdef cl_int err
        cdef cl_event *c_wait
        
        num_events = create_events(event_list, &c_wait)
        
        if c_wait == NULL:
            raise ValueError("Wait event list cannot be empty")
        try:
            with nogil:
                err = clEnqueueWaitForEvents(self.id, num_events, c_wait)
        finally:
            PyMem_Free(c_wait)
        check_error("Enqueueing Wait for Events", err)
        
    def enqueue_barrier(self):
        cdef cl_int err
        
        with nogil:
            err = clEnqueueBarrier(self.id)
        
        check_error("Enqueueing Barrier", err)
    def finish(self):
        cdef cl_int err
        
        with nogil:
            err = clFinish(self.id)
        check_error("Finish", err)
    
    def flush(self):
        cdef cl_int err
        
        with nogil:
            err = clFlush(self.id)
        check_error("Flush", err)

cdef class Buffer(MemObject):
    pass

cdef class Image2D(Image):
    def __cinit__(self):
        self.dim = 2

cdef class Image3D(Image):
    def __cinit__(self):
        self.dim = 3

cdef class Event(Referenced):
    types = {
        CL_EVENT_COMMAND_QUEUE: CommandQueue,
        CL_EVENT_COMMAND_TYPE: c_command_type,
        CL_EVENT_COMMAND_EXECUTION_STATUS: c_command_status,
        CL_EVENT_REFERENCE_COUNT: "cl_uint",
        CL_PROFILING_COMMAND_QUEUED: "cl_ulong",
        CL_PROFILING_COMMAND_SUBMIT: "cl_ulong",
        CL_PROFILING_COMMAND_START: "cl_ulong",
        CL_PROFILING_COMMAND_END: "cl_ulong",
    }
    def __cinit__(self):
        self.info = <InfoFunc> clGetEventInfo
        self._retain = <ReferenceFunc> clRetainEvent
        self._release = <ReferenceFunc> clReleaseEvent
        self.ref_count_param = CL_EVENT_REFERENCE_COUNT
        self.params = c_event_info
        self.param_types = self.types   
        
    cpdef wait(self):
        cdef cl_int err
        cdef cl_event event = self.id
        with nogil:
            err = clWaitForEvents(1, &event)
        check_error("Waiting on event", err)
        
    def get_profiling_info(self, Constant param_name):
        return self.c_get_info(param_name.value, <InfoFunc> clGetEventProfilingInfo, self.param_types, NULL)

cdef class Program(Referenced):
    types = {
        CL_PROGRAM_REFERENCE_COUNT: 'cl_uint',
        CL_PROGRAM_CONTEXT: Context,
        CL_PROGRAM_NUM_DEVICES: 'cl_uint',
        CL_PROGRAM_DEVICES: (Device,),
        CL_PROGRAM_BINARY_SIZES: 'size_t[]',
        CL_PROGRAM_BUILD_STATUS: c_build_status,
    }
    def __cinit__(self):
        self.info = <InfoFunc> clGetProgramInfo
        self._retain = <ReferenceFunc> clRetainProgram
        self._release = <ReferenceFunc> clReleaseProgram
        self.ref_count_param = CL_PROGRAM_REFERENCE_COUNT
        self.params = c_program_info
        self.param_types = self.types
        
    def build(self, device_list = None, options = None,suppress_error = False):
        cdef char* c_opts = NULL
        cdef cl_device_id *c_devices = NULL
        cdef cl_uint device_count = 0
        cdef unsigned int i
        cdef Device e
        cdef cl_int err
        
        if options is not None:
            c_opts = options
        try:
            if device_list is not None:
                device_count = len(device_list)
                c_devices = <cl_device_id *> PyMem_Malloc(sizeof(cl_device_id) * device_count)
                for i,e in enumerate(device_list):
                    c_devices[i] = e.id
            # Don't use user callback.  This call is blocking
            with nogil:
                err = clBuildProgram(self.id, device_count, c_devices, c_opts, NULL, NULL)
        finally:
            PyMem_Free(c_devices)
        if suppress_error and err == CL_BUILD_PROGRAM_FAILURE:
            return False
        check_error("Build Program",err)
        return True
            
    def get_build_info(self, Device device, Constant param_name):
        return self.c_get_info(param_name.value, <InfoFunc>clGetProgramBuildInfo, self.param_types, device.id)

    def create_kernel(self, kernel_name):
        cdef bytes n = str(kernel_name)
        cdef char *c_name = n
        cdef cl_kernel id
        cdef cl_int err
        
        id = clCreateKernel(self.id, c_name, &err)
        check_error('Create Kernel', err)
        return referenced_factory(Kernel, id, self)
        
cdef class Kernel(Referenced):
    types = {
        CL_KERNEL_NUM_ARGS: 'cl_uint',
        CL_KERNEL_REFERENCE_COUNT: 'cl_uint',
        CL_KERNEL_CONTEXT: Context,
        CL_KERNEL_PROGRAM: Program,
        CL_KERNEL_WORK_GROUP_SIZE: 'size_t',
        CL_KERNEL_COMPILE_WORK_GROUP_SIZE: 'size_t[]',
        CL_KERNEL_LOCAL_MEM_SIZE: 'cl_ulong',
    }
    def __cinit__(self):
        self.info = <InfoFunc> clGetKernelInfo
        self._retain = <ReferenceFunc> clRetainKernel
        self._release = <ReferenceFunc> clReleaseKernel
        self.ref_count_param = CL_KERNEL_REFERENCE_COUNT
        self.params = c_kernel_info
        self.param_types = self.types
    
    def get_work_group_info(self,Device device, Constant param_name):        
        return self.c_get_info(param_name.value, <InfoFunc> clGetKernelWorkGroupInfo, self.param_types, device.id)
  
    def set_arg(self, int arg_index, arg_value):
        cdef Py_buffer view
        cdef void *ptr = NULL
        cdef size_t size
        view.buf = NULL
        
        if arg_value is None:
            ptr = NULL
            size = 0
        elif isinstance(arg_value, MemObject):
            ptr = &((<MemObject>arg_value).id)
            size = sizeof(cl_mem)
        elif isinstance(arg_value, Sampler):
            ptr = &((<Sampler>arg_value).id)
            size = sizeof(cl_sampler)
        elif isinstance(arg_value, LocalMemory):
            ptr = NULL
            size = (<LocalMemory>arg_value).size
        else:
            #use buffer interface.  
            #Works with strings made by struct or one element ndarrays or ctypes
            size = get_buffer(arg_value, False, &ptr, &view)
        try:
            check_error("Setting Kernel Argument",
                        clSetKernelArg(self.id, arg_index, size, ptr))
        finally:
            release_buffer(&view)
      
      
cdef class Sampler(Referenced):
    types = {
        CL_SAMPLER_REFERENCE_COUNT : 'cl_uint',
        CL_SAMPLER_CONTEXT : Context,
        CL_SAMPLER_ADDRESSING_MODE : c_addressing_mode,
        CL_SAMPLER_FILTER_MODE : c_filter_mode,
        CL_SAMPLER_NORMALIZED_COORDS : 'cl_bool',
    }

    def __cinit__(self):
        self.info = <InfoFunc> clGetSamplerInfo
        self._retain = <ReferenceFunc> clRetainSampler
        self._release = <ReferenceFunc> clReleaseSampler
        self.ref_count_param = CL_SAMPLER_REFERENCE_COUNT
        self.params = c_sampler_info
        self.param_types = self.types    
#-------------------------------------------------------------------
# global functions
#-----------------------------------------------------------------------
cpdef wait_for_events(event_list):
    cdef cl_uint num_events
    cdef cl_event * c_wait
    cdef cl_int err
    
    num_events = create_events(event_list, &c_wait)
    
    try:
        with nogil:
            err = clWaitForEvents(num_events, c_wait)
    finally:
        PyMem_Free(c_wait)
    check_error("Wait For Events", err)

cpdef Context create_context(properties=None, devices = None):
    cdef cl_bitfield d_type
    cdef cl_int err
    cdef cl_context_properties * c_props = NULL
    cdef cl_device_id * ids = NULL
    cdef size_t size 
    cdef unsigned int i = 0
    cdef cl_context_properties prop
    cdef Device d
    cdef cl_context id
    cdef Constant key
    
    if devices is None:
        raise ValueError("Devices cannot be None")
    
    #makes it simpler later
    if isinstance(devices, Device):
        devices = (devices,)
    
    try:
        if (properties is not None):
            size = len(properties)
            #two items for every property plus 1 for the null termination
            c_props = <cl_context_properties*> PyMem_Malloc(sizeof(cl_context_properties) * size * 2+ 1)
        
            for key,value in properties.iteritems():
                c_props[i] = <cl_context_properties> key.value
                if isinstance(value, InfoGetter):               
                    prop = <cl_context_properties> (<InfoGetter?> value).id
                elif isinstance(value, Constant):
                    prop = <cl_context_properties> (<Constant?> value).value
                else:
                    prop = <cl_context_properties> int(value)
                c_props[i+1] = prop
                i+=2
            c_props[i] = <cl_context_properties> NULL

        size = len(devices)
        ids = <cl_device_id *> PyMem_Malloc(sizeof(cl_device_id) * size)
        for i in range(size):
            ids[i] = (<Device> devices[i]).id
        id = clCreateContext(c_props, size, ids, NULL, NULL, &err)
        check_error("Create Context", err)
    finally:
        PyMem_Free(ids)    
        PyMem_Free(c_props)
    
    return referenced_factory(Context, id, None)
    

cpdef Context create_context_from_type(properties = None, device_type = CL_DEVICE_TYPE_DEFAULT,
        Platform platform = None):

    cdef cl_context_properties * c_props = NULL
    cdef size_t size 
    cdef unsigned int i = 0
    cdef cl_context_properties prop
    cdef cl_platform_id plat_id
    cdef cl_int err
    cdef cl_device_type d_type = device_type
    cdef cl_context id
    
    if platform is None:
        plat_id = get_first_platform()
    else:
        plat_id = platform.id
    
    try:
        if (properties is None):
            size = 0
        else:
            size = len(properties)
            
        #two items for every property plus 1 for the null termination plus 2 for platform spec
        c_props = <cl_context_properties*> PyMem_Malloc(sizeof(cl_context_properties) * size * 2+ 3)
    
        if properties is not None:
            for key,value in properties.iteritems():
                constant = key
                c_props[i] = <cl_context_properties> constant.value
                if isinstance(value, InfoGetter):               
                    prop = <cl_context_properties> (<InfoGetter?> value).id
                elif isinstance(value, Constant):
                    prop = <cl_context_properties> (<Constant?> value).value
                else:
                    prop = <cl_context_properties> int(value)
                c_props[i+1] = prop
                i+=2

        c_props[i] = <cl_context_properties> CL_CONTEXT_PLATFORM
        c_props[i+1] = <cl_context_properties> plat_id
        c_props[i+2] = <cl_context_properties> NULL
            
        id = clCreateContextFromType(c_props, d_type, NULL, NULL, &err)
        
    finally:
        PyMem_Free(c_props)
    check_error("Create Context", err)    
    return referenced_factory(Context, id, None)
    


cpdef get_platforms():
    cdef cl_platform_id *ids = NULL
    cdef cl_uint num
    cdef int x
    cdef Platform p
    
    check_error("Get Platforms", clGetPlatformIDs(0, NULL, &num))
    ids = <cl_platform_id *> PyMem_Malloc(sizeof(cl_platform_id) * num)
    try:
        check_error("Get Platforms", clGetPlatformIDs(num, ids, NULL))

        result = []
        for x in range(num):
            p = info_getter_factory(Platform, ids[x])
            result.append(p)
        
    finally:
        PyMem_Free(ids)
    
    return result
